Conversation
…ing; defaults must obey declared knob domains (fixes NVIDIA#653) Two changes that are one rule: a scheduler the served kernel cannot decode must be impossible to reach — whether requested or defaulted. 1. NVIDIA#585 port (fixes NVIDIA#653): the LPT/LPT_L2 tile schedulers landed in prefill_d128_fp8_sm100.py but the hunk-symmetric port to the SM107 sibling was never done, so the adapter's auto causal sched policy picked SCHED_LPT_L2 and every causal/masked per-tensor FP8 graph on cc10.7 failed plan build (29 of 33 suite cases). Ported: qh_per_kh/seqlen_kv threaded through the five warp groups and appended at all ten scheduler decode call sites, mirroring the SM100 file position for position. 2. Domain honesty end to end: - Both fp8 kernel files declare SUPPORTED_SCHED_POLICIES (decode truth at the source). - The adapter gains _sm100_sched_domain(); check_support now admits explicit NATURAL/LPT/LPT_L2 requests (was NATURAL-only — under- declared), and the causal defaulting heuristic CLAMPS its pick into the served route's domain, falling back to NATURAL. A default may never be a value the route could not honor as an explicit request. - The three SM100-family engine rows declare the true request domain {NATURAL, LPT, LPT_L2} (previously {NATURAL} while the lowering auto-picked LPT — dishonest in both directions). SM120 rows untouched. - frost/README knob-channel section documents the rule. 3. Sibling-lockstep fences (test_sdpa_fp8_sibling_parity.py): source-scan tests (no GPU/DSL needed) assert every decode call site carries the LPT args and the sched-domain constants match; module tests assert compile() signatures in lockstep (NVIDIA#574's fence), kernel domains == adapter table, and CFG divergence is exactly the intentional Rubin set {TILE_K_HW_BMM1, TILE_K_HW_BMM2, STAGES_KV}. Negative-tested: removing one call site's args turns the fence red. SM100 box: 40 passed (5 parity + 2 routing + 33 fp8 e2e). Rubin board validation of the causal suite to follow on the PR. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
vedaanta
added
cat-bug
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (6)
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review. 📝 WalkthroughWalkthroughSM100 and SM107 FP8 SDPA routes now support matching natural, LPT, and LPT_L2 scheduler domains. The SM107 kernel receives required decode parameters across its execution paths. GPU-independent tests enforce sibling parity and routing consistency. ChangesFP8 scheduler parity
Estimated code review effort: 4 (Complex) | ~45 minutes Merge Risk: ⚪ Minimal · up to The PR ports the scheduler behavior, aligns declared domains and defaults, and adds parity coverage; no actionable merge-blocking risk remains beyond normal checks and review. Possibly related PRs
Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
|
Needs conflict resolution/CI etc. |
YangXu1990uiuc
left a comment
YangXu1990uiuc
left a comment
There was a problem hiding this comment.
Codex bot review · model: gpt-6-astra
Reviewed f418fd8251f6 against base c61a5e4eb5c5.
The SM107 FP8 decode-argument port fixes the reported LPT_L2 compilation failure, independently reproduced on Rubin. The broader domain declaration still needs a THD restriction, and the source-only tests should avoid importing the optional DSL adapter. No P0 found.
[P2] Restrict the new scheduler domain for THD plans (source)
This row also serves THD, so the expanded domain now accepts explicit LPT/LPT_L2 requests for ragged plans. But _dispatch_decode_initial/_payload in _common_sm100.py return _thd_decode whenever _thd_on is true; neither requested dense scheduler runs. check_support and compile both succeed, silently ignoring the requested scheduling policy. Make the effective domain depend on THD as well as the route, and reject unsupported non-NATURAL requests in both the engine probe and standalone adapter unless the THD scheduler actually implements them. The parity test only checks the two dense FP8 siblings and misses this combination.
Evidence: SM100 direct-API probe with B=2,H=4,S=256,D=128: the baseline rejects THD policies 1/2, while the PR accepts and compiles both. Source tracing confirms both THD dispatch functions bypass _decode_initial/_payload unconditionally. Attribution: Newly reachable relative to branch point 8a758b8: this PR expands both the row domain and adapter gate while leaving the THD decoder unchanged. Develop has also broadened scheduler eligibility independently.
[P2] Locate source files without importing the DSL adapter (source)
The first two tests intentionally have no requires_dsl marker, but _kernel_dir imports api_dsl, which imports APIBase and then cutlass. In a core-only environment these new L0 tests fail instead of remaining source-only; the other three tests correctly skip. Derive the kernel directory from the lightweight cudnn package location (or the test checkout) without importing the adapter, and update the sibling paths for the current kernel-directory layout when rebasing.
Evidence: Dependency-isolation run with CUTLASS unavailable: 2 failed, 3 skipped. Both failures trace through _kernel_dir -> api_dsl -> api_base -> import cutlass; with DSL installed all five tests pass. Attribution: Introduced by the new unmarked source tests and their location helper.
Validation: Built a matching extension from the reviewed head. SM100: 38 parity/FP8 tests passed. Rubin baseline 8a758b8: 29 failed, 4 passed, with the failures reporting missing qh_per_kh/seqlen_kv for SCHED_LPT_L2. Rubin head: 38 passed. Both runs explicitly selected L0 and L1. Checked all 10 initial/payload decode calls in the modified Rubin sibling and the corresponding SM100 calls. Direct-API probes covered dense/THD, causal/noncausal and NATURAL/LPT/LPT_L2, separately from the reference suite. Additional Rubin persistent-handoff probe: B=4,Hq=8,Hkv=2,S=2048,D=128, causal GQA with KV lengths [2048,1536,1024,768]. The 512-CTA grid exceeds the GPU's 208 SMs. Both LPT and LPT_L2 passed the FP32-dequant reference and produced bit-identical outputs; the actual compiled scheduler policy was verified.
Limitations: Current develop has substantial integration conflicts and per-flavor scheduler restrictions that this older branch predates. Preserve its distinction between an unspecified policy and explicit NATURAL; the PR head still substitutes a causal policy for explicit NATURAL, which also happens on its baseline and is already addressed in current develop. The tentative merge pairs the old SM107 d128 file with kernels/sm107/prefill_d192_d128_fp8.py. Reapply the port to the actual d128 route deliberately, preserve current per-flavor restrictions and THD default handling, and update test paths; I have not validated a resolved current-develop integration. The previously failing Rubin LPT_L2 path has no valid before-fix timing baseline. No end-to-end performance claim is made here. There is no pending fork test workflow or test pipeline in the PR thread; after resolving the conflicts, request fresh CI with a standalone @cudnn-ci-bot run frost comment.
Approved under the trial's no-P0 threshold; any P1/P2 findings above remain for the owner to address. Merge timing stays with the owner.
3 participants
Fixes #653.
Two changes that are one rule: a scheduler the served kernel cannot decode must be unreachable — whether requested or defaulted.
1. The #585 port (the bug fix)
#585 landed the LPT/LPT_L2 tile schedulers in
prefill_d128_fp8_sm100.py, but the hunk-symmetric port to the SM107 sibling was never done. The adapter's auto causal sched policy picksSCHED_LPT_L2, the sibling's decode guard refuses, and every causal/masked per-tensor FP8 graph on cc10.7 failed plan build (29 of 33 suite cases). Ported position-for-position with the SM100 file:qh_per_kh/seqlen_kvthreaded through the five warp groups and appended at all ten scheduler decode call sites.2. Domain honesty end to end
SUPPORTED_SCHED_POLICIES— decode truth at the source._sm100_sched_domain();check_supportnow admits explicit NATURAL/LPT/LPT_L2 requests (was NATURAL-only — under-declared while the lowering auto-picked LPT), and the causal defaulting heuristic clamps its pick into the served route's domain, falling back to NATURAL. A default may never be a value the route could not honor as an explicit request — the invariant that turns this bug class into a clean fallback.{NATURAL, LPT, LPT_L2}(previously{NATURAL}— dishonest in both directions). SM120 rows untouched.frost/READMEknob-channel section documents the rule.3. Sibling-lockstep fences (
test_sdpa_fp8_sibling_parity.py)Third silent-drift incident for this kernel pair (#574 signature skew, #602 amax ABI, #585/#653). The fences: source-scan tests (no GPU/DSL — run in any CI) assert every decode call site carries the LPT args and the sched-domain constants match; module tests assert
compile()signatures in lockstep, kernel domains == adapter table, and CFG divergence is exactly the intentional Rubin set{TILE_K_HW_BMM1, TILE_K_HW_BMM2, STAGES_KV}. Negative-tested: removing one call site's args turns the fence red.Testing
🤖 Generated with Claude Code
Summary by CodeRabbit
New Features
Bug Fixes
Tests